
#pragma once

#include <stdexcept>
#include <string>
#include <stdio.h>
#include <math.h>
#include "kernel_defs.hpp"
#include "maths_functions.hpp"

template<gpu_size_t BLOCK_SIZE, typename Complex, typename Real>
__global__
void phase_shift_kernel(const Complex* d_in,
                        Complex*       d_out,
                        Real           coef,
                        Real           delay_rate,
                        gpu_size_t     count,
                        Real           fringe,
                        Real           dd_fringe,
                        gpu_size_t    fft_size,
                        size_t        fft_cnt,
                        float    update_time)
{
  gpu_size_t thread_idx = threadIdx.x + blockIdx.x*BLOCK_SIZE;
  gpu_size_t thread_count = BLOCK_SIZE * gridDim.x;

    // Iterate over values in the input
  for( gpu_size_t idx=thread_idx; idx<count; idx+=thread_count ) {
    Complex a = d_in[idx];
    const Real pi = 3.14159265358979323846f;
    size_t seq = 0;
    Real coeff_mod;
    seq = (int) idx / fft_size;
  
    //calculating the actual phase shift in radians from the fstc,fringe and rate of change of fringe  
    coeff_mod =((idx%fft_size)*(coef - (delay_rate*update_time*(seq+fft_cnt)))) + ((fringe*pi/180.f) - (dd_fringe*update_time*(pi/180.f)*(seq+fft_cnt)));

   Complex phase;
    // TODO: Check if pre-computing these and loading from mem is faster
    __sincosf(coeff_mod, &phase.y, &phase.x);

    Complex result;
    result.x = a.x*phase.x + a.y*phase.y;
    result.y = a.y*phase.x - a.x*phase.y;
    d_out[idx] = result;
  }
}



template<typename Complex, typename Real>
void phase_shift_device(const Complex* d_in,
                        Complex*       d_out,
                        gpu_size_t     count,
                        Real           frac,
                        Real          delay_rate,                       
                        Real          fringe,
                        Real          dd_fringe,
                        float    update_time,
                        gpu_size_t    fft_size,
                        size_t        fft_cnt,
                        cudaStream_t   stream = 0)
{
  enum { BLOCK_SIZE = 128  };
  dim3 block(BLOCK_SIZE);
  dim3 grid(math::min(math::div_round_up(count,
                                         (gpu_size_t)BLOCK_SIZE),
                      (gpu_size_t)CUDA_MAX_GRID_SIZE));

  const float pi = 3.14159265358979323846f;
  Real coef = pi * frac / fft_size;
  delay_rate = pi * delay_rate / fft_size;
  phase_shift_kernel<BLOCK_SIZE><<<grid, block, 0, stream>>>(d_in,
                                                             d_out,
                                                             coef,
                                                             delay_rate,
                                                             count,
                                                             fringe,
                                                             dd_fringe,
                                                             fft_size,
                                                             fft_cnt,
                                                             update_time);
#ifdef _DEBUG
  cudaStreamSynchronize(stream);
  cudaError_t error = cudaGetLastError();
  if( error != cudaSuccess ) {
    throw std::runtime_error(std::string("In phase_shift_kernel: ") +
                             cudaGetErrorString(error));
  }
#endif
}

